#include "kernel_operator.h"
#include "kernel_pow_scalar.h"
#include "kernel_pow_tensor.h"
#include "kernel_pow_broadcast.h"

extern "C" __global__ __aicore__ void pows(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
{
    TPipe pipe;
    GET_TILING_DATA(tiling_data, tiling);
    if (TILING_KEY_IS(3)){
        KernelBroadCast op;
        op.Init(&pipe, x1, x2, y,
                tiling_data.y_dimensional,
                tiling_data.y_ndarray, tiling_data.x1_ndarray, tiling_data.x2_ndarray,
                tiling_data.y_sumndarray, tiling_data.x1_sumndarray, tiling_data.x2_sumndarray);
        op.Process();
    }else if (TILING_KEY_IS(1)){
        KernelOptensor<DTYPE_X1> op;
        op.Init(&pipe, x1, x2, y,
                tiling_data.coreDataNum, tiling_data.finalTileNum, tiling_data.tileDataNum, tiling_data.tailDataNum);
        op.Process();
    }else if (TILING_KEY_IS(2)){
        KernelOpscalar<DTYPE_X1> op;
        op.Init(&pipe, x1, x2, y,
                tiling_data.coreDataNum, tiling_data.finalTileNum, tiling_data.tileDataNum, tiling_data.tailDataNum);
        op.Process();
    }
}
